
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include "../cuda-memory-debug/cuda-memory-debug-defs.h"


__global__ void WrongSize(int * x)
{
    *x = 22;
}

__global__ void OverwriteEnd(int * x)
{
    *x = 1;
    *(x + 1) = 2;
}

__global__ void OverwriteStart(int * x)
{
    *x = 3;
    *(x - 1) = 4;
}

// Out of the Ocelot memory checker, an example ...

#include <string>

__global__ void init(int* data)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    data[tid] = tid;
}

void wrongMemcpyDirection()
{
    int* hostA = new int[128];
    memset(hostA, 0, sizeof(int)*128);
    int* deviceA;
    if (cudaMalloc( (void**) &deviceA, sizeof(int)*128 ))
        printf("hi\n");
    cudaMemcpy( hostA, deviceA, sizeof(int)*128, cudaMemcpyHostToDevice );
    cudaFree( deviceA );
    delete[] hostA;
}

void hostPointerOnDevice()
{
    int* hostA = new int[128];
    memset(hostA, 0, sizeof(int)*128);
    int* deviceA;
    cudaMalloc( (void**) &deviceA, sizeof(int)*128 );
    cudaMemcpy( deviceA, hostA, sizeof(int)*128, cudaMemcpyHostToDevice );
    init<<< 4, 32 >>>(hostA);
    cudaFree( deviceA );
    delete[] hostA;
    cudaThreadExit();
}

void bufferOverrun()
{
    int* hostA = new int[128];
    memset(hostA, 0, sizeof(int)*128);
    int* deviceA;
    cudaMalloc( (void**) &deviceA, sizeof(int)*128 );
    cudaMemcpy( deviceA, hostA, sizeof(int)*128, cudaMemcpyHostToDevice );
    init<<< 4, 33 >>>( deviceA );
    cudaFree( deviceA );
    delete[] hostA; 
}


int main()
{
    int cuda_device = 0;
    cudaError_t errSetDevice = cudaSetDevice(cuda_device);
    if (errSetDevice != 0)
    {
        std::cout << "Error setting device " << errSetDevice << std::endl;
        return 1;
    }

    cudaMalloc((void**)123456789, 1);
    dim3 Dg(1, 1);      // Number of blocks
    dim3 Db(1);         // Size of blocks
    size_t Ns = 0;

    //CudaMemoryDebug::SetTraceAllCalls(true, __FILE__, __LINE__);


    ////////////////////////////////////
    // Test overwrites.
    ////////////////////////////////////

    {
        std::cout << "Clean Test...\n";
        // Memory overwrite due to cudaMemcpy.
        short * d;
        cudaError_t e1 = cudaMalloc((void**)&d, sizeof(short));  // Not correct size.
        short h = 5;
        int errMemcpy1 = cudaMemcpy(d, (void*)&h, sizeof(short), cudaMemcpyHostToDevice);
        if (errMemcpy1 != 0)
        {
            std::cout << "Error in cuda memcpy " << errMemcpy1 << std::endl;
        }
        cudaFree(d);
        std::cout << std::endl;
    }
    {
        std::cout << "Test 1...\n";
        // Memory overwrite due to cudaMemcpy.
        int * d = 0;
        cudaError_t e1 = cudaMalloc((void**)&d, sizeof(short));  // Not correct size.
        int h = 5;
        int errMemcpy1 = cudaMemcpy(d, (void*)&h, sizeof(int), cudaMemcpyHostToDevice);
        if (errMemcpy1 != 0)
        {
            std::cout << "Error in cuda memcpy " << errMemcpy1 << std::endl;
        }
        cudaFree(d);
        std::cout << std::endl;
        cudaError_t resetErrKernel = cudaGetLastError();
    }

    {
        std::cout << "Test 2...\n";
        // Memory overwrite due to kernel.
        short * d = 0;
        cudaError_t e1 = cudaMalloc((void**)&d, sizeof(short));  // Not correct size.
        short h = 5;
        int errMemcpy1 = cudaMemcpy(d, (void*)&h, sizeof(short), cudaMemcpyHostToDevice);
        if (errMemcpy1 != 0)
        {
            std::cout << "Error in cuda memcpy " << errMemcpy1 << std::endl;
        }
        WrongSize<<<Dg, Db, Ns>>>((int*)d);
        cudaThreadSynchronize();
        cudaError_t errKernel = cudaGetLastError();
        if (errKernel)
        {
            std::cout << "Error in kernel call " << errKernel << std::endl;
        }
        cudaFree(d);
        std::cout << std::endl;
        cudaError_t resetErrKernel = cudaGetLastError();
    }
    {
        std::cout << "Test 3...\n";
        // Memory overwrite due to kernel.
        short * d;
        cudaError_t e1 = cudaMalloc((void**)&d, sizeof(short));  // Not correct size.
        int h = 5;
        int errMemcpy1 = cudaMemcpy(d, (void*)&h, sizeof(short), cudaMemcpyHostToDevice);
        if (errMemcpy1 != 0)
        {
            std::cout << "Error in cuda memcpy " << errMemcpy1 << std::endl;
        }
        OverwriteEnd<<<Dg, Db, Ns>>>((int*)d);
        cudaThreadSynchronize();
        cudaError_t errKernel = cudaGetLastError();
        if (errKernel)
        {
            std::cout << "Error in kernel call " << errKernel << std::endl;
        }
        cudaFree(d);
        std::cout << std::endl;
        cudaError_t resetErrKernel = cudaGetLastError();
    }
    {
        std::cout << "Test 4...\n";
        // Memory overwrite due to kernel.
        short * d;
        cudaError_t e1 = cudaMalloc((void**)&d, sizeof(short));  // Not correct size.
        int h = 5;
        int errMemcpy1 = cudaMemcpy(d, (void*)&h, sizeof(short), cudaMemcpyHostToDevice);
        if (errMemcpy1 != 0)
        {
            std::cout << "Error in cuda memcpy " << errMemcpy1 << std::endl;
        }
        OverwriteStart<<<Dg, Db, Ns>>>((int*)d);
        cudaThreadSynchronize();
        cudaError_t errKernel = cudaGetLastError();
        if (errKernel)
        {
            std::cout << "Error in kernel call " << errKernel << std::endl;
        }
        cudaFree(d);
        std::cout << std::endl;
        cudaError_t resetErrKernel = cudaGetLastError();
    }

    {
        std::cout << "Test 5...\n";
        // Huge memory allocation via cudaHostAlloc.
        struct cudaDeviceProp prop;
        cudaError_t e = cudaGetDeviceProperties(&prop, cuda_device);
        int size = prop.totalGlobalMem;
        if (prop.canMapHostMemory)
        {
            unsigned char * big;
            int size = 100000000;
            int errMalloc = cudaHostAlloc((void**)&big, size, cudaHostAllocMapped);
            int * dbig;
            int err = cudaHostGetDevicePointer((void**)&dbig, (void*)big, 0);
            short * d;
            OverwriteStart<<<Dg, Db, Ns>>>((int*)dbig);
            cudaThreadSynchronize();
            cudaError_t errKernel = cudaGetLastError();
            if (errKernel)
            {
                std::cout << "Error in kernel call " << errKernel << std::endl;
            }
            cudaFreeHost(big);
        }
        std::cout << std::endl;
        cudaError_t resetErrKernel = cudaGetLastError();
    }
    {
        wrongMemcpyDirection();
        hostPointerOnDevice();
        bufferOverrun();
    }

    ////////////////////////////////////
    //  Allocation buffers to leak.
    ////////////////////////////////////
    int * d2;
    cudaError_t e2 = cudaMalloc((void**)&d2, sizeof(int));

    return 0;
}

